Adds benchmarks for cub::DeviceMerge
#3529
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
elstehle
force-pushed
the
enh/cub-merge-benchmarks
branch
from
16568b1 to
6bb30c6
Compare
🟩 CI finished in 1h 05m: Pass: 100%/90 | Total: 15h 30m | Avg: 10m 20s | Max: 51m 03s | Hits: 414%/12772
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental | |
| python | |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental | |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 90)
| # | Runner |
|---|---|
| 65 | linux-amd64-cpu16 |
| 11 | linux-amd64-gpu-v100-latest-1 |
| 9 | windows-amd64-cpu16 |
| 4 | linux-arm64-cpu16 |
| 1 | linux-amd64-gpu-h100-latest-1-testing |
🟨 CI finished in 3h 02m: Pass: 98%/90 | Total: 21h 02m | Avg: 14m 01s | Max: 1h 00m | Hits: 422%/10928
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental | |
| python | |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental | |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 90)
| # | Runner |
|---|---|
| 65 | linux-amd64-cpu16 |
| 11 | linux-amd64-gpu-v100-latest-1 |
| 9 | windows-amd64-cpu16 |
| 4 | linux-arm64-cpu16 |
| 1 | linux-amd64-gpu-h100-latest-1-testing |
🟨 CI finished in 1d 01h: Pass: 98%/89 | Total: 15h 03m | Avg: 10m 09s | Max: 42m 25s | Hits: 422%/10928
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental | |
| python | |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental | |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 89)
| # | Runner |
|---|---|
| 65 | linux-amd64-cpu16 |
| 11 | linux-amd64-gpu-v100-latest-1 |
| 8 | windows-amd64-cpu16 |
| 4 | linux-arm64-cpu16 |
| 1 | linux-amd64-gpu-h100-latest-1-testing |
🟩 CI finished in 4h 15m: Pass: 100%/89 | Total: 1d 16h | Avg: 27m 07s | Max: 1h 15m | Hits: 236%/10936
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental | |
| python | |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental | |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 89)
| # | Runner |
|---|---|
| 65 | linux-amd64-cpu16 |
| 11 | linux-amd64-gpu-v100-latest-1 |
| 8 | windows-amd64-cpu16 |
| 4 | linux-arm64-cpu16 |
| 1 | linux-amd64-gpu-h100-latest-1 |
🟩 CI finished in 1h 57m: Pass: 100%/89 | Total: 16h 36m | Avg: 11m 11s | Max: 1h 48m | Hits: 422%/10936
|
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental | |
| python | |
| CCCL C Parallel Library | |
| Catch2Helper |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental | |
| +/- | python |
| +/- | CCCL C Parallel Library |
| +/- | Catch2Helper |
🏃 Runner counts (total jobs: 89)
| # | Runner |
|---|---|
| 65 | linux-amd64-cpu16 |
| 8 | windows-amd64-cpu16 |
| 6 | linux-amd64-gpu-rtxa6000-latest-1 |
| 4 | linux-arm64-cpu16 |
| 3 | linux-amd64-gpu-rtx4090-latest-1 |
| 2 | linux-amd64-gpu-rtx2080-latest-1 |
| 1 | linux-amd64-gpu-h100-latest-1 |
gevtushenko
approved these changes
Feb 1, 2025
Comment on lines
+72
to
+108
| const auto num_items_lhs = elements / 2; | ||
| const auto num_items_rhs = elements - num_items_lhs; | ||
| auto counting_it = thrust::make_counting_iterator(offset_t{0}); | ||
| thrust::copy_if( | ||
| counting_it, | ||
| counting_it + elements, | ||
| rnd_selector_val.begin(), | ||
| thrust::make_tabulate_output_iterator(write_pivot_point_t<offset_t>{ | ||
| static_cast<offset_t>(num_items_lhs), thrust::raw_pointer_cast(pivot_point.data())}), | ||
| select_lhs_op); | ||
|
|
||
| thrust::device_vector<key_t> keys_lhs(num_items_lhs); | ||
| thrust::device_vector<key_t> keys_rhs(num_items_rhs); | ||
| thrust::device_vector<key_t> keys_out(elements); | ||
|
|
||
| // Generate increasing input range to sample from | ||
| thrust::device_vector<key_t> increasing_input = generate(elements); | ||
| thrust::sort(increasing_input.begin(), increasing_input.end()); | ||
|
|
||
| // Select lhs from input up to pivot point | ||
| offset_t pivot_point_val = pivot_point[0]; | ||
| auto const end_lhs = thrust::copy_if( | ||
| increasing_input.cbegin(), | ||
| increasing_input.cbegin() + pivot_point_val, | ||
| rnd_selector_val.cbegin(), | ||
| keys_lhs.begin(), | ||
| select_lhs_op); | ||
|
|
||
| // Select rhs items from input up to pivot point | ||
| auto const end_rhs = thrust::copy_if( | ||
| increasing_input.cbegin(), | ||
| increasing_input.cbegin() + pivot_point_val, | ||
| rnd_selector_val.cbegin(), | ||
| keys_rhs.begin(), | ||
| select_rhs_op); | ||
| // From pivot point copy all remaining items to rhs | ||
| thrust::copy(increasing_input.cbegin() + pivot_point_val, increasing_input.cbegin() + elements, end_rhs); |
There was a problem hiding this comment.
remark: this is beautiful 💚
Comment on lines
+45
to
+111
| // Retrieve axis parameters | ||
| const auto elements = static_cast<std::size_t>(state.get_int64("Elements{io}")); | ||
| const bit_entropy entropy = str_to_entropy(state.get_string("Entropy")); | ||
|
|
||
| // We generate data distributions in the range [0, 255] that, with lower entropy, get skewed towards 0. | ||
| // We use this to generate increasingly large *consecutive* segments of data that are getting selected from the lhs | ||
| thrust::device_vector<uint8_t> rnd_selector_val = generate(elements, entropy); | ||
| uint8_t threshold = 128; | ||
| select_if_less_than_t select_lhs_op{false, threshold}; | ||
| select_if_less_than_t select_rhs_op{true, threshold}; | ||
|
|
||
| // The following algorithm only works under the precondition that there's at least 50% of the data in the lhs | ||
| // If that's not the case, we simply swap the logic for selecting into lhs and rhs | ||
| const auto num_items_selected_into_lhs = | ||
| static_cast<offset_t>(thrust::count_if(rnd_selector_val.begin(), rnd_selector_val.end(), select_lhs_op)); | ||
| if (num_items_selected_into_lhs < elements / 2) | ||
| { | ||
| using ::cuda::std::swap; | ||
| swap(select_lhs_op, select_rhs_op); | ||
| } | ||
|
|
||
| // We want lhs and rhs to be of equal size. We also want to have skewed distributions, such that we put different | ||
| // workloads on the binary search part. For this reason, we identify the index from the input, referred to as pivot | ||
| // point, after which the lhs is "full". We compose the rhs by selecting all items up to the pivot point that were not | ||
| // selected for lhs and *all* items after the pivot point. | ||
| constexpr std::size_t num_pivot_points = 1; | ||
| thrust::device_vector<offset_t> pivot_point(num_pivot_points); | ||
| const auto num_items_lhs = elements / 2; | ||
| const auto num_items_rhs = elements - num_items_lhs; | ||
| auto counting_it = thrust::make_counting_iterator(offset_t{0}); | ||
| thrust::copy_if( | ||
| counting_it, | ||
| counting_it + elements, | ||
| rnd_selector_val.begin(), | ||
| thrust::make_tabulate_output_iterator(write_pivot_point_t<offset_t>{ | ||
| static_cast<offset_t>(num_items_lhs), thrust::raw_pointer_cast(pivot_point.data())}), | ||
| select_lhs_op); | ||
|
|
||
| thrust::device_vector<key_t> keys_lhs(num_items_lhs); | ||
| thrust::device_vector<key_t> keys_rhs(num_items_rhs); | ||
| thrust::device_vector<key_t> keys_out(elements); | ||
| thrust::device_vector<value_t> values_lhs(num_items_lhs); | ||
| thrust::device_vector<value_t> values_rhs(num_items_rhs); | ||
| thrust::device_vector<value_t> values_out(elements); | ||
|
|
||
| // Generate increasing input range to sample from | ||
| thrust::device_vector<key_t> increasing_input = generate(elements); | ||
| thrust::sort(increasing_input.begin(), increasing_input.end()); | ||
|
|
||
| // Select lhs from input up to pivot point | ||
| offset_t pivot_point_val = pivot_point[0]; | ||
| auto const end_lhs = thrust::copy_if( | ||
| increasing_input.cbegin(), | ||
| increasing_input.cbegin() + pivot_point_val, | ||
| rnd_selector_val.cbegin(), | ||
| keys_lhs.begin(), | ||
| select_lhs_op); | ||
|
|
||
| // Select rhs items from input up to pivot point | ||
| auto const end_rhs = thrust::copy_if( | ||
| increasing_input.cbegin(), | ||
| increasing_input.cbegin() + pivot_point_val, | ||
| rnd_selector_val.cbegin(), | ||
| keys_rhs.begin(), | ||
| select_rhs_op); | ||
| // From pivot point copy all remaining items to rhs | ||
| thrust::copy(increasing_input.cbegin() + pivot_point_val, increasing_input.cbegin() + elements, end_rhs); |
There was a problem hiding this comment.
suggestion: this looks like a duplicate of workload generation you added in the keys version. Maybe we could simplify the code a bit by having cuda::std::pair<thrust::device_vector<T>, thrust::device_vector<T>> generate(elements, entropy); auto [keys_lhs, keys_rhs] = ... function.
There was a problem hiding this comment.
Sounds good! I'll do that in a follow-up PR 👍
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Description
Closes #3528